#include "kernel_operator.h"
using namespace AscendC;
constexpr int32_t BUFFER_NUM = 1; 
constexpr float A[] = {
    (float)(4.65128586073990045278E-5), (float)(7.31589045238094711071E-3),
    (float)(1.33847639578309018650E-1), (float)(8.79691311754530315341E-1),
    (float)(2.71149851196553469920E0),  (float)(4.25697156008121755724E0),
    (float)(3.29771340985225106936E0),  (float)(1.00000000000000000126E0),
};

constexpr float B[] = {
    (float)(6.90990488912553276999E-4), (float)(2.54043763932544379113E-2),
    (float)(2.82974860602568089943E-1), (float)(1.41172597751831069617E0),
    (float)(3.63800533345137075418E0),  (float)(5.03278880143316990390E0),
    (float)(3.54771340985225096217E0),  (float)(9.99999999999999998740E-1),
};

class KernelSpence {
public:
    __aicore__ inline KernelSpence() {}
    __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, uint32_t magic, uint32_t totalLength, uint32_t blockLength,
                              uint32_t tileNum, uint32_t tileLength,
                              uint32_t lasttileLength, uint32_t tilingKey)
    {
        ASSERT(GetBlockNum() != 0 && "block dim can not be zero!");
        this->blockLength = blockLength;
        this->tileNum =
            tileNum ASSERT(tileNum != 0 && "tile num can not be zero!");
        this->tileLength = tileLength / BUFFER_NUM;
        this->lasttileLength = lasttileLength;
        this->totalLength = totalLength;
        this->magic = magic;
        this->key = tilingKey;

        xGm.SetGlobalBuffer((__gm__ DTYPE_X*)x + this->blockLength * GetBlockIdx(),
                            this->blockLength);
        yGm.SetGlobalBuffer((__gm__ DTYPE_X*)y + this->blockLength * GetBlockIdx(),
                            this->blockLength);
        

        this->pi = (float)3.14159265358979323846;//3.1415926f; //3.14159265358979323846;
        this->PI2O6 = (float)(pi * pi) / 6 ;


        pipe.InitBuffer(inQueueIN, BUFFER_NUM, 64 * sizeof(DTYPE_X));
        pipe.InitBuffer(outQueueOUT, BUFFER_NUM, 64 * sizeof(DTYPE_X));

        pipe.InitBuffer(tempBuf1, 32 * sizeof(float));
        pipe.InitBuffer(tempBuf2, 32 * sizeof(float));
        pipe.InitBuffer(tempBuf3, 32 * sizeof(float));
        pipe.InitBuffer(tempBuf4, 32 * sizeof(float));
        pipe.InitBuffer(tempBuf5, 32 * sizeof(float));

        pipe.InitBuffer(zero, 32 * sizeof(float));
        pipe.InitBuffer(one, 32 * sizeof(float));
        pipe.InitBuffer(two, 32 * sizeof(float));
        pipe.InitBuffer(three_halves, 32 * sizeof(float));
        pipe.InitBuffer(half_1, 32 * sizeof(float));

        pipe.InitBuffer(cmp1, 32 * sizeof(int16_t));
        pipe.InitBuffer(cmp2, 32 * sizeof(int16_t));
        pipe.InitBuffer(cmp3, 32 * sizeof(int16_t));
        pipe.InitBuffer(cmp4, 32 * sizeof(int16_t));
        pipe.InitBuffer(cmp, 32 * sizeof(int16_t));
        pipe.InitBuffer(cmp_8, 32 * sizeof(uint8_t));
        pipe.InitBuffer(cmp_f, 32 * sizeof(half));

        pipe.InitBuffer(x1, 32 * sizeof(float));
        pipe.InitBuffer(x2, 32 * sizeof(float));

        pipe.InitBuffer(conv_x1, 32 * sizeof(float));
        pipe.InitBuffer(conv_y1, 32 * sizeof(float));

    }
    __aicore__ inline void Process()
    {
        int32_t loopCount = this->tileNum;
        for (int32_t i = 0; i < loopCount; i++) {
            CopyIn(i);
            Compute(i, (DTYPE_X)0);
            CopyOut(i);
        }
    }

private:
    __aicore__ inline void CopyIn(int32_t progress)
    {
        LocalTensor<DTYPE_X> inLocal = inQueueIN.AllocTensor<DTYPE_X>();
        DataCopy(inLocal[0], xGm[progress * this->tileLength], this->tileLength * 2);
        inQueueIN.EnQue(inLocal);
    }
    __aicore__ inline void Compute(int32_t progress, float flag)
    {
        LocalTensor<float> inLocal = inQueueIN.DeQue<float>();
        LocalTensor<float> xLocal = inLocal;
        LocalTensor<float> yLocal = outQueueOUT.AllocTensor<float>();

        LocalTensor<float> temp1 = tempBuf1.Get<float>();
        LocalTensor<float> temp2 = tempBuf2.Get<float>();
        LocalTensor<float> temp3 = tempBuf3.Get<float>();
        LocalTensor<float> temp4 = tempBuf4.Get<float>();
        LocalTensor<float> temp5 = tempBuf5.Get<float>();
        
        LocalTensor<float> zero_ = zero.Get<float>();
        LocalTensor<float> one_ = one.Get<float>();
        LocalTensor<float> two_ = two.Get<float>();
        LocalTensor<float> three_halves_ = three_halves.Get<float>();
        LocalTensor<float> half_1_ = half_1.Get<float>();

        LocalTensor<int16_t> cmp_ = cmp.Get<int16_t>();
        LocalTensor<int16_t> cmp_1 = cmp1.Get<int16_t>();
        LocalTensor<int16_t> cmp_2 = cmp2.Get<int16_t>();
        LocalTensor<int16_t> cmp_3 = cmp3.Get<int16_t>();
        LocalTensor<int16_t> cmp_4 = cmp4.Get<int16_t>();
        LocalTensor<uint8_t> cmp_8_ = cmp_8.Get<uint8_t>();
        LocalTensor<half> cmp_f_ = cmp_f.Get<half>();

        LocalTensor<float> x_1 = x1.Get<float>();
        LocalTensor<float> x_2 = x2.Get<float>();
        // LocalTensor<float> spence_ = spence.Get<float>();
        
        if (progress == 0) {
            Duplicate(zero_, (float)0, this->tileLength);
            Duplicate(one_, (float)1, this->tileLength);
            Duplicate(two_, (float)2, this->tileLength);
            Duplicate(three_halves_, (float)1.5, this->tileLength);
            Duplicate(half_1_, (float)0.5, this->tileLength);
        }

        Duplicate(temp1, (float)0, this->tileLength);
        Duplicate(temp2, (float)0, this->tileLength);
        Duplicate(temp3, (float)0, this->tileLength);
        Duplicate(temp4, (float)0, this->tileLength);
        Duplicate(temp5, (float)0, this->tileLength);

        Duplicate(cmp_, (int16_t)0, this->tileLength);
        Duplicate(cmp_1, (int16_t)0, this->tileLength);
        Duplicate(cmp_2, (int16_t)0, this->tileLength);
        Duplicate(cmp_3, (int16_t)0, this->tileLength);
        Duplicate(cmp_4, (int16_t)0, this->tileLength);
        
        Duplicate(x_1, (float)0, this->tileLength);
        Duplicate(x_2, (float)0, this->tileLength);


        Compare(cmp_8_, xLocal, zero_, CMPMODE::NE, this->tileLength);
        Select(yLocal, cmp_8_, xLocal, this->PI2O6, SELMODE::VSEL_TENSOR_SCALAR_MODE, this->tileLength);
        Cast(cmp_f_, cmp_8_, RoundMode::CAST_NONE, this->tileLength);
        Cast(cmp_1, cmp_f_, RoundMode::CAST_RINT, this->tileLength);

        Compare(cmp_8_, xLocal, one_, CMPMODE::NE, this->tileLength);
        Select(yLocal, cmp_8_, yLocal, (float)0, SELMODE::VSEL_TENSOR_SCALAR_MODE, this->tileLength);
        Cast(cmp_f_, cmp_8_, RoundMode::CAST_NONE, this->tileLength);
        Cast(cmp_2, cmp_f_, RoundMode::CAST_RINT, this->tileLength);


        And(cmp_, cmp_1, cmp_2, this->tileLength);


        // Reciprocal(temp1, xLocal, this->tileLength);
        Div(temp1, one_, xLocal, this->tileLength);
        Compare(cmp_8_, xLocal, two_, CMPMODE::LT, this->tileLength);
        Select(x_1, cmp_8_, xLocal, temp1, SELMODE::VSEL_CMPMASK_SPR, this->tileLength);
        Cast(cmp_f_, cmp_8_, RoundMode::CAST_NONE, this->tileLength);
        Cast(cmp_1, cmp_f_, RoundMode::CAST_RINT, this->tileLength);


        Compare(cmp_8_, three_halves_, x_1,  CMPMODE::LT, this->tileLength);
        Cast(cmp_f_, cmp_8_, RoundMode::CAST_NONE, this->tileLength);
        Cast(cmp_1, cmp_f_, RoundMode::CAST_RINT, this->tileLength);

        Compare(cmp_8_, x_1, half_1_, CMPMODE::LT, this->tileLength);
        Cast(cmp_f_, cmp_8_, RoundMode::CAST_NONE, this->tileLength);
        Cast(cmp_2, cmp_f_, RoundMode::CAST_RINT, this->tileLength);
        
        Div(temp1, one_, x_1, this->tileLength);
        // Reciprocal(temp1, x_1, this->tileLength);
        Sub(temp1, temp1, one_, this->tileLength);
        Muls(temp2, x_1, (float)-1, this->tileLength);
        Sub(temp3, x_1, one_, this->tileLength);

        Not(cmp_3, cmp_1, this->tileLength);
        Not(cmp_4, cmp_2, this->tileLength);
        And(cmp_2, cmp_3, cmp_2, this->tileLength);
        And(cmp_4, cmp_3, cmp_4, this->tileLength);

        Select(x_2, cmp_1, temp1, x_2, SELMODE::VSEL_CMPMASK_SPR, this->tileLength);
        Select(x_2, cmp_2, temp2, x_2, SELMODE::VSEL_CMPMASK_SPR, this->tileLength);
        Select(x_2, cmp_4, temp3, x_2, SELMODE::VSEL_CMPMASK_SPR, this->tileLength);

        Duplicate<float>(temp4, (float)0, this->tileLength);
        Duplicate<float>(temp5, (float)0, this->tileLength);

        for(int i = 0; i < 8; i++) {
            Mul(temp4, temp4, x_2, this->tileLength);
            Mul(temp5, temp5, x_2, this->tileLength);
            Adds(temp4, temp4, A[i], this->tileLength);
            Adds(temp5, temp5, B[i], this->tileLength);
        }

        Div(temp1, temp4, temp5, this->tileLength);
        Mul(temp1, temp1, x_2, this->tileLength);
        Muls(yLocal, temp1, (float)-1.0, this->tileLength);

        Sub(temp1, one_, x_1, this->tileLength);
        Ln(temp1, temp1, this->tileLength);
        Ln(temp3, x_1, this->tileLength);

        Mul(temp2, temp3, temp1, this->tileLength);
        Muls(temp2, temp2, (float)-1.0, this->tileLength);
        Adds(temp2, temp2, this->PI2O6, this->tileLength);
        Sub(temp2, temp2, yLocal, this->tileLength);

        Compare(cmp_8_, x_1, half_1_, CMPMODE::LT, this->tileLength);
        Cast(cmp_f_, cmp_8_, RoundMode::CAST_NONE, this->tileLength);
        Cast(cmp_1, cmp_f_, RoundMode::CAST_RINT, this->tileLength);
        And(cmp_1, cmp_1, cmp_, this->tileLength);
        Select(yLocal, cmp_1, temp2, yLocal, SELMODE::VSEL_CMPMASK_SPR, this->tileLength);

        Mul(temp1, temp3, temp3, this->tileLength);
        Muls(temp2, half_1_, (float)-1.0, this->tileLength);
        Mul(temp2, temp1, temp2, this->tileLength);
        Sub(temp2, temp2, yLocal, this->tileLength);

        Compare(cmp_8_, three_halves_, xLocal, CMPMODE::LT, this->tileLength);
        Cast(cmp_f_, cmp_8_, RoundMode::CAST_NONE, this->tileLength);
        Cast(cmp_2, cmp_f_, RoundMode::CAST_RINT, this->tileLength);
        And(cmp_2, cmp_, cmp_2, this->tileLength);
        Select(yLocal, cmp_2, temp2, yLocal, SELMODE::VSEL_CMPMASK_SPR, this->tileLength);

        Compare(cmp_8_, xLocal, zero_, CMPMODE::NE, this->tileLength);
        Select(yLocal, cmp_8_, yLocal, this->PI2O6, SELMODE::VSEL_TENSOR_SCALAR_MODE, this->tileLength);

        Compare(cmp_8_, xLocal, one_, CMPMODE::NE, this->tileLength);
        Select(yLocal, cmp_8_, yLocal, (float)0, SELMODE::VSEL_TENSOR_SCALAR_MODE, this->tileLength);

        outQueueOUT.EnQue<float>(yLocal);
        inQueueIN.FreeTensor(inLocal);
    }

    __aicore__ inline void Compute(int32_t progress, half flag)
    {
        LocalTensor<half> inLocal = inQueueIN.DeQue<half>();
        LocalTensor<half> xLocal = inLocal;
        LocalTensor<half> yLocal = outQueueOUT.AllocTensor<half>();

        LocalTensor<float> xLocal_32 = conv_x1.Get<float>();
        LocalTensor<float> yLocal_32 = conv_y1.Get<float>();

        LocalTensor<float> temp1 = tempBuf1.Get<float>();
        LocalTensor<float> temp2 = tempBuf2.Get<float>();
        LocalTensor<float> temp3 = tempBuf3.Get<float>();
        LocalTensor<float> temp4 = tempBuf4.Get<float>();
        LocalTensor<float> temp5 = tempBuf5.Get<float>();
        
        LocalTensor<float> zero_ = zero.Get<float>();
        LocalTensor<float> one_ = one.Get<float>();
        LocalTensor<float> two_ = two.Get<float>();
        LocalTensor<float> three_halves_ = three_halves.Get<float>();
        LocalTensor<float> half_1_ = half_1.Get<float>();

        LocalTensor<int16_t> cmp_ = cmp.Get<int16_t>();
        LocalTensor<int16_t> cmp_1 = cmp1.Get<int16_t>();
        LocalTensor<int16_t> cmp_2 = cmp2.Get<int16_t>();
        LocalTensor<int16_t> cmp_3 = cmp3.Get<int16_t>();
        LocalTensor<int16_t> cmp_4 = cmp4.Get<int16_t>();
        LocalTensor<uint8_t> cmp_8_ = cmp_8.Get<uint8_t>();
        LocalTensor<half> cmp_f_ = cmp_f.Get<half>();

        LocalTensor<float> x_1 = x1.Get<float>();
        LocalTensor<float> x_2 = x2.Get<float>();
        // LocalTensor<float> spence_ = spence.Get<float>();
        
        Cast(xLocal_32, xLocal, RoundMode::CAST_NONE, this->tileLength);

        if (progress == 0) {
            Duplicate(zero_, (float)0, this->tileLength);
            Duplicate(one_, (float)1, this->tileLength);
            Duplicate(two_, (float)2, this->tileLength);
            Duplicate(three_halves_, (float)1.5, this->tileLength);
            Duplicate(half_1_, (float)0.5, this->tileLength);
        }

        Duplicate(temp1, (float)0, this->tileLength);
        Duplicate(temp2, (float)0, this->tileLength);
        Duplicate(temp3, (float)0, this->tileLength);
        Duplicate(temp4, (float)0, this->tileLength);
        Duplicate(temp5, (float)0, this->tileLength);

        Duplicate(cmp_, (int16_t)0, this->tileLength);
        Duplicate(cmp_1, (int16_t)0, this->tileLength);
        Duplicate(cmp_2, (int16_t)0, this->tileLength);
        Duplicate(cmp_3, (int16_t)0, this->tileLength);
        Duplicate(cmp_4, (int16_t)0, this->tileLength);
        
        Duplicate(x_1, (float)0, this->tileLength);
        Duplicate(x_2, (float)0, this->tileLength);


        Compare(cmp_8_, xLocal_32, zero_, CMPMODE::NE, this->tileLength);
        Select(yLocal_32, cmp_8_, xLocal_32, this->PI2O6, SELMODE::VSEL_TENSOR_SCALAR_MODE, this->tileLength);
        Cast(cmp_f_, cmp_8_, RoundMode::CAST_NONE, this->tileLength);
        Cast(cmp_1, cmp_f_, RoundMode::CAST_RINT, this->tileLength);

        Compare(cmp_8_, xLocal_32, one_, CMPMODE::NE, this->tileLength);
        Select(yLocal_32, cmp_8_, yLocal_32, (float)0, SELMODE::VSEL_TENSOR_SCALAR_MODE, this->tileLength);
        Cast(cmp_f_, cmp_8_, RoundMode::CAST_NONE, this->tileLength);
        Cast(cmp_2, cmp_f_, RoundMode::CAST_RINT, this->tileLength);


        And(cmp_, cmp_1, cmp_2, this->tileLength);


        // Reciprocal(temp1, xLocal, this->tileLength);
        Div(temp1, one_, xLocal_32, this->tileLength);
        Compare(cmp_8_, xLocal_32, two_, CMPMODE::LT, this->tileLength);
        Select(x_1, cmp_8_, xLocal_32, temp1, SELMODE::VSEL_CMPMASK_SPR, this->tileLength);
        Cast(cmp_f_, cmp_8_, RoundMode::CAST_NONE, this->tileLength);
        Cast(cmp_1, cmp_f_, RoundMode::CAST_RINT, this->tileLength);


        Compare(cmp_8_, three_halves_, x_1,  CMPMODE::LT, this->tileLength);
        Cast(cmp_f_, cmp_8_, RoundMode::CAST_NONE, this->tileLength);
        Cast(cmp_1, cmp_f_, RoundMode::CAST_RINT, this->tileLength);

        Compare(cmp_8_, x_1, half_1_, CMPMODE::LT, this->tileLength);
        Cast(cmp_f_, cmp_8_, RoundMode::CAST_NONE, this->tileLength);
        Cast(cmp_2, cmp_f_, RoundMode::CAST_RINT, this->tileLength);
        
        Div(temp1, one_, x_1, this->tileLength);
        // Reciprocal(temp1, x_1, this->tileLength);
        Sub(temp1, temp1, one_, this->tileLength);
        Muls(temp2, x_1, (float)-1, this->tileLength);
        Sub(temp3, x_1, one_, this->tileLength);

        Not(cmp_3, cmp_1, this->tileLength);
        Not(cmp_4, cmp_2, this->tileLength);
        And(cmp_2, cmp_3, cmp_2, this->tileLength);
        And(cmp_4, cmp_3, cmp_4, this->tileLength);

        Select(x_2, cmp_1, temp1, x_2, SELMODE::VSEL_CMPMASK_SPR, this->tileLength);
        Select(x_2, cmp_2, temp2, x_2, SELMODE::VSEL_CMPMASK_SPR, this->tileLength);
        Select(x_2, cmp_4, temp3, x_2, SELMODE::VSEL_CMPMASK_SPR, this->tileLength);


        Duplicate<float>(temp1, (float)1.0, this->tileLength);
        Duplicate<float>(temp2, (float)1.0, this->tileLength);
        Duplicate<float>(temp3, (float)1.0, this->tileLength);
        Duplicate<float>(temp4, (float)A[7], this->tileLength);
        Duplicate<float>(temp5, (float)B[7], this->tileLength);


        for (int i = 6; i >= 0; i--) {
            Mul(temp1, temp1, x_2, this->tileLength);
            Muls(temp2, temp1, B[i], this->tileLength);
            Muls(temp3, temp1, A[i], this->tileLength);
            Add(temp4, temp4, temp3, this->tileLength);
            Add(temp5, temp5, temp2, this->tileLength);
        }

        Div(temp1, temp4, temp5, this->tileLength);
        Mul(temp1, temp1, x_2, this->tileLength);
        Muls(yLocal_32, temp1, (float)-1.0, this->tileLength);

        Sub(temp1, one_, x_1, this->tileLength);
        Ln(temp1, temp1, this->tileLength);
        Ln(temp3, x_1, this->tileLength);

        Mul(temp2, temp3, temp1, this->tileLength);
        Muls(temp2, temp2, (float)-1.0, this->tileLength);
        Adds(temp2, temp2, this->PI2O6, this->tileLength);
        Sub(temp2, temp2, yLocal_32, this->tileLength);

        Compare(cmp_8_, x_1, half_1_, CMPMODE::LT, this->tileLength);
        Cast(cmp_f_, cmp_8_, RoundMode::CAST_NONE, this->tileLength);
        Cast(cmp_1, cmp_f_, RoundMode::CAST_RINT, this->tileLength);
        And(cmp_1, cmp_1, cmp_, this->tileLength);
        Select(yLocal_32, cmp_1, temp2, yLocal_32, SELMODE::VSEL_CMPMASK_SPR, this->tileLength);

        Mul(temp1, temp3, temp3, this->tileLength);
        Muls(temp2, half_1_, (float)-1.0, this->tileLength);
        Mul(temp2, temp1, temp2, this->tileLength);
        Sub(temp2, temp2, yLocal_32, this->tileLength);

        Compare(cmp_8_, three_halves_, xLocal_32, CMPMODE::LT, this->tileLength);
        Cast(cmp_f_, cmp_8_, RoundMode::CAST_NONE, this->tileLength);
        Cast(cmp_2, cmp_f_, RoundMode::CAST_RINT, this->tileLength);
        And(cmp_2, cmp_, cmp_2, this->tileLength);
        Select(yLocal_32, cmp_2, temp2, yLocal_32, SELMODE::VSEL_CMPMASK_SPR, this->tileLength);

        Compare(cmp_8_, xLocal_32, zero_, CMPMODE::NE, this->tileLength);
        Select(yLocal_32, cmp_8_, yLocal_32, this->PI2O6, SELMODE::VSEL_TENSOR_SCALAR_MODE, this->tileLength);

        Compare(cmp_8_, xLocal_32, one_, CMPMODE::NE, this->tileLength);
        Select(yLocal_32, cmp_8_, yLocal_32, (float)0, SELMODE::VSEL_TENSOR_SCALAR_MODE, this->tileLength);

        Cast(yLocal, yLocal_32, RoundMode::CAST_NONE, this->tileLength);

        outQueueOUT.EnQue<half>(yLocal);
        inQueueIN.FreeTensor(inLocal);
    }


    __aicore__ inline void CopyOut(int32_t progress)
    {
        LocalTensor<DTYPE_Y> outLocal = outQueueOUT.DeQue<DTYPE_Y>();
        DataCopy(yGm[progress * this->tileLength], outLocal, this->tileLength * 2);
        outQueueOUT.FreeTensor(outLocal);
    }

private:
    TPipe pipe;
    TQue<QuePosition::VECIN, BUFFER_NUM> inQueueIN;
    TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueOUT;
    GlobalTensor<DTYPE_X> xGm;
    GlobalTensor<DTYPE_Y> yGm;

    int key;

    float pi;
    TBuf<> tempBuf1;
    TBuf<> tempBuf2;
    TBuf<> tempBuf3;
    TBuf<> tempBuf4;
    TBuf<> tempBuf5;
    TBuf<> zero;
    TBuf<> one;
    TBuf<> two;
    TBuf<> three_halves;
    TBuf<> half_1;
    TBuf<> cmp1;
    TBuf<> cmp2;
    TBuf<> cmp3;
    TBuf<> cmp4;
    TBuf<> cmp;
    TBuf<> cmp_8;
    TBuf<> cmp_f;

    TBuf<> x1;
    TBuf<> x2;

    TBuf<> conv_x1;
    TBuf<> conv_y1;

    uint32_t magic;
    uint32_t totalLength;
    uint32_t blockLength;
    uint32_t tileNum;
    uint32_t tileLength;
    uint32_t lasttileLength;

    // pi**2 / 6.
    float PI2O6;

};


extern "C" __global__ __aicore__ void spence(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling) {
    GET_TILING_DATA(tiling_data, tiling);
    // TODO: user kernel impl
    KernelSpence op;
    uint32_t tilingKey = 1;
    if (TILING_KEY_IS(1)) {
        tilingKey = 1;
    } else if (TILING_KEY_IS(2)) {
        tilingKey = 2;
    } else {
        tilingKey = 1;
    }
    op.Init(x, y, tiling_data.magic, tiling_data.totalLength, 
          tiling_data.blockLength,  
          tiling_data.tileNum, tiling_data.tileLength,
          tiling_data.lasttileLength, tilingKey);
    op.Process();
}